
#include "kernel_operator.h"

#include "heaviside_sca.h"
#include "heaviside_spec1.h"

extern "C" __global__ __aicore__ void heaviside(GM_ADDR input, GM_ADDR values, GM_ADDR out,
                                                GM_ADDR workspace, GM_ADDR tiling) {
    GET_TILING_DATA(tiling_data, tiling);
    // TODO: user kernel impl
    TPipe pipe;

    if (TILING_KEY_IS(1)) {
        HeavisideSpec1 spec1;
        spec1.Init(input, values, out, tiling_data.inputSize, tiling_data.valueSize, &pipe);
        spec1.Process();
    } else if (TILING_KEY_IS(2)) {
        HeavisideSca sca;
        sca.Init(input, values, out, tiling_data.inputSize, tiling_data.valueSize, &pipe);
        sca.Process();
    }
}